/*******************************************

  TRABALHO DE CONCORRENTES - FINAL

  Smooth - Paralelizado(Cuda)

********************************************/

/*********************
Bibliotecas
*********************/
#include <stdio.h>
#include <stdlib.h>
#include <math.h>

/*********************
Definições
*********************/


/*********************
Funções gerais
*********************/
__global__ void smooth( int *dev_X, int *dev_nX, int tamLin, int tamCol);

/******************
Função Principal
*******************/
int main (void){
	
	FILE *pfile;
	char firstLine[3];
	int lines, columns, maxPixel;
	int sizeInt;

	int **R; // Matriz que estará armazenada a imagem - Cor R
	int **G; // Matriz que estará armazenada a imagem - Cor G
	int **B; // Matriz que estará armazenada a imagem - Cor B
	int *nR; // nova Matriz que estará armazenada a imagem - Cor R
	int *nG; // nova Matriz que estará armazenada a imagem - Cor G
	int *nB; // nova Matriz que estará armazenada a imagem - Cor B

	int *dev_R; 
	int *dev_G; 
	int *dev_B; 
	int *dev_nR; 
	int *dev_nG; 
	int *dev_nB; 

	int i, j, k=0;

	/**********************************
	Leitura das informações iniciais
	**********************************/
	scanf("%s", firstLine);			// Le a primeira linha "P3"
	scanf("%d", &lines);			// Le a qntLinhas
	scanf("%d", &columns);			// Le a qntColunas
	scanf("%d", &maxPixel);			// Le o valor max do Pixels

	/**********************************
	Aloca as matrizes definindo um range de 2 pixel a mais como borda
	**********************************/
	R = (int **) calloc((lines+4),sizeof(int *));
	G = (int **) calloc((lines+4),sizeof(int *));
	B = (int **) calloc((lines+4),sizeof(int *));
	nR = (int *) calloc((lines+4)*(columns+4),sizeof(int));
	nG = (int *) calloc((lines+4)*(columns+4),sizeof(int));
	nB = (int *) calloc((lines+4)*(columns+4),sizeof(int));
	for(i=0;i<(lines+4);i++)
	{
		R[i] = (int *) calloc((columns+4),sizeof(int));
		G[i] = (int *) calloc((columns+4),sizeof(int));
		B[i] = (int *) calloc((columns+4),sizeof(int));
	}

	/**********************************
	Leitura dos valores da matriz e configura as bordas
	**********************************/
	for(i = 0; i<(lines+4); i++){
		for(j = 0; j < (columns+4); j++)
			if(i<=1 || i>=lines+2)
			{
				R[i][j]=0;
				G[i][j]=0;
				B[i][j]=0;
			}
			else if(j<=1 || j>=columns+2)
			{
				R[i][j]=0;
				G[i][j]=0;
				B[i][j]=0;
			} else
				scanf("%d %d %d ", &R[i][j], &G[i][j], &B[i][j]);
	}

	// vetoriza matriz
	for(i=0;i<(lines+4);i++){
		for(j=0;j<(columns+4);j++)
		{
			nR[k] = R[i][j];
			nG[k] = G[i][j];
			nB[k] = B[i][j];
			k++;
		}
	}

	/**********************************
	Copiando informações para a Cuda
	**********************************/
	sizeInt = (lines+4) * (columns+4) * sizeof(int);

	// Alocando as matrizes Cuda
	cudaMalloc((void**) &dev_R, sizeInt);
	cudaMalloc((void**) &dev_G, sizeInt);
	cudaMalloc((void**) &dev_B, sizeInt);
	cudaMalloc((void**) &dev_nR, sizeInt);
	cudaMalloc((void**) &dev_nG, sizeInt);
	cudaMalloc((void**) &dev_nB, sizeInt);

	//printf("Alfa R: %d %d %d %d %d SizeInt%d\n", R[0][0], R[0][1], R[0][2], R[0][3], R[0][4], sizeInt);
	// printf("Alfa nR: %d %d %d %d %d\n", nR[0], nR[1], nR[2], nR[3], nR[4]);
	// Copiando do Host para o Device
	cudaMemcpy(dev_R, nR, sizeInt, cudaMemcpyHostToDevice);
	cudaMemcpy(dev_G, nG, sizeInt, cudaMemcpyHostToDevice);
	cudaMemcpy(dev_B, nB, sizeInt, cudaMemcpyHostToDevice);

	// Invocação do Kernel
	dim3 dimBlock(4, 4);
	dim3 dimGrid((int)ceil(lines/dimBlock.x),(int)ceil(columns/dimBlock.y));
	smooth<<<dimGrid,dimBlock>>>(dev_R, dev_nR, lines+4, columns+4); 
	smooth<<<dimGrid,dimBlock>>>(dev_G, dev_nG, lines+4, columns+4);
	smooth<<<dimGrid,dimBlock>>>(dev_B, dev_nB, lines+4, columns+4);

	//Copia o Cuda para a Memória
	cudaMemcpy(nR, dev_nR, sizeInt, cudaMemcpyDeviceToHost);
	cudaMemcpy(nG, dev_nG, sizeInt, cudaMemcpyDeviceToHost);
	cudaMemcpy(nB, dev_nB, sizeInt, cudaMemcpyDeviceToHost);

	//printf("Alfa R: %d %d %d %d %d\n", R[0][0], R[0][1], R[0][2], R[0][3], R[0][4]);
	//printf("Alfa nR: %d %d %d %d %d %d\n", nR[0], nR[1], nR[2], nR[3], nR[4], nR[5]);
	
	cudaFree(dev_R); cudaFree(dev_B); cudaFree(dev_G);
	cudaFree(dev_nR); cudaFree(dev_nB); cudaFree(dev_nG);

	for(i = 0; i<(lines+4); i++)
		for(j = 0; j < (columns+4); j++)
		{
			R[i][j] = nR[((i*(columns+4))+j)];
			G[i][j] = nG[((i*(columns+4))+j)];
			B[i][j] = nB[((i*(columns+4))+j)];
		}

	/**********************************
	Salvando imagem no arquivo
	**********************************/
	pfile = fopen ("saidaCuda.ppm", "w");
	//printf("teste\n");
	fprintf(pfile, "%s \n", firstLine);
	fprintf(pfile, "%d ", (lines));
	fprintf(pfile, "%d \n", (columns));
	fprintf(pfile, "%d \n", maxPixel);


	for(i = 2; i<lines+2;i++)
		for(j = 2; j<columns+2; j++)	
				if((R[i][j]>=0 && R[i][j]<=maxPixel) && (G[i][j]>=0 && G[i][j]<=maxPixel) && (B[i][j]>=0 && B[i][j]<=maxPixel))
					fprintf(pfile, "%d %d %d\n", R[i][j], G[i][j], B[i][j]);
				else
					fprintf(pfile, "0 0 0\n", R[i][j], G[i][j], B[i][j]);

	/**********************************
	Liberando Matrizes
	**********************************/
	for(i=0;i<(lines+4);i++)
	{
		free(R[i]); free(G[i]); free(B[i]);
	}

	free(R); free(G); free(B);
	free(nR); free(nG); free(nB);

	return 0;
}


__global__ void smooth( int *dev_X, int *dev_nX, int tamLin, int tamCol) {
	
	int col = blockIdx.x*blockDim.x+threadIdx.x;
	int row = blockIdx.y*blockDim.y+threadIdx.y;

	int index = col + row * tamCol;

	if ((col < tamCol) && (row <= tamLin)) {	

		//if(row >= 2 && row <= (tamLin+2)){
		//	if((col >= 2) && (col <=tamCol+2))
		//	{
				dev_nX[index] =
					  dev_X[((col-2)+(row-2)*tamCol)] +
					  dev_X[((col-1)+(row-2)*tamCol)] +
					  dev_X[((col)+(row-2)*tamCol)] +
					  dev_X[((col+1)+(row-2)*tamCol)] +
					  dev_X[((col+2)+(row-2)*tamCol)] +

					  dev_X[((col-2)+(row-1)*tamCol)] +
					  dev_X[((col-1)+(row-1)*tamCol)] +
					  dev_X[((col)+(row-1)*tamCol)] +
					  dev_X[((col+1)+(row-1)*tamCol)] +
					  dev_X[((col+2)+(row-1)*tamCol)] +

					  dev_X[((col-2)+(row)*tamCol)] +
					  dev_X[((col-1)+(row)*tamCol)] +
					  dev_X[((col)+(row)*tamCol)] +
					  dev_X[((col+1)+(row)*tamCol)] +
					  dev_X[((col+2)+(row)*tamCol)] +

					  dev_X[((col-2)+(row+1)*tamCol)] +
					  dev_X[((col-1)+(row+1)*tamCol)] +
					  dev_X[((col)+(row+1)*tamCol)] +
					  dev_X[((col+1)+(row+1)*tamCol)] +
					  dev_X[((col+2)+(row+1)*tamCol)] +

					  dev_X[((col-2)+(row+2)*tamCol)] +
					  dev_X[((col-1)+(row+2)*tamCol)] +
					  dev_X[((col)+(row+2)*tamCol)] +
					  dev_X[((col+1)+(row+2)*tamCol)] +
					  dev_X[((col+2)+(row+2)*tamCol)];

				dev_nX[index] /= 25;
		/*	}
			else
			{
				dev_nX[index] = 0;
			}
		}*/
	} else dev_nX[index] = 0;
	
 
}
